#include "Common.h"

#define CUDA_CHECK(call) \
	if((call) != cudaSuccess) { \
        	cudaError_t err = cudaGetLastError(); \
        	fprintf(stderr, "RCUDA error calling \"#call\", code is %d\n", err); \
		exit(0);\
	}

typedef struct para{
	int rw, rh, iw, ih, rk, ts;
	color *rc, *ic;
	ACCM *pl, *cl;
}PARA;

static color *refc, *inpc;
static ACCM  *preline, *curline;
static int buf_size;

static __device__ int IPX[][2] = {{-1,0},{-1,-1},{-1,1},{-2,0},{-2,-1},{-2,1},{0,0}};
static __device__ int IPY[][2] = {{0,-1},{-1,-1},{1,-1},{0,-2},{-1,-2},{1,-2},{0,0}};

static color* copyImageToDevice(Image *image);
static __device__ short calDistance(color *c1, color *c2);
static __device__ ACCM * get(ACCM *, int, int, int);

__global__ void accms(int m, int n, PARA p, int size);
__device__ PARA *paras;

extern "C" void gpu_initializeDeviceBuf(int s){
	refc = copyImageToDevice(&ref);
	inpc = copyImageToDevice(&inp);
	buf_size = s;
	CUDA_CHECK(cudaMalloc((void**)&preline, sizeof(ACCM) * inp.width * inp.height * buf_size));
	CUDA_CHECK(cudaMalloc((void**)&curline, sizeof(ACCM) * inp.width * inp.height * buf_size));
}

//2.
extern "C" void gpu_initCurStart(ACCM *value){
	CUDA_CHECK(cudaMemcpy(&curline[0], value, sizei, cudaMemcpyHostToDevice));
}

//3.
extern "C" void gpu_accmcalc(int m, int n, int st) {
	PARA p;
	p.rw = ref.width;
	p.rh = ref.height;
	p.iw = inp.width;
	p.ih = inp.height;
	p.rk = rank + 1;
	p.ts = st;
	p.rc = refc;
	p.ic = inpc;
	p.pl = preline;
	p.cl = curline;
	int size = inp.width * inp.height;
	accms<<<BLOCK_NUM, THREAD_NUM>>>(m, n, p, size);
}

//4.
extern "C" void gpu_copyResult(ACCM *buf, int ri){
	CUDA_CHECK(cudaMemcpy(buf, &curline[ri*inp.width*inp.height], sizei, cudaMemcpyDeviceToHost));
}

//6.
extern "C" void gpu_switchPreAndCurLine(){
	ACCM* tmp = preline;
	preline = curline;
	curline = tmp;
	CUDA_CHECK(cudaMemset(curline, 0, sizeof(ACCM) * buf_size * inp.width * inp.height));
}

extern "C" void gpu_cudaFinalize(){
	CUDA_CHECK(cudaFree(refc));
	CUDA_CHECK(cudaFree(inpc));
	CUDA_CHECK(cudaFree(preline));
	CUDA_CHECK(cudaFree(curline));
}

__global__ void accms(int m, int n, PARA p, int size){
	int i, j, x, y, l, distance;
	int MIN, path, cur;

	paras = &p;
	cur = blockIdx.x * THREAD_NUM + threadIdx.x;

	while(cur < size){
		i = cur / paras->ih;
		j = cur % paras->ih;
		distance = calDistance(&paras->rc[(paras->ts + m) * paras->rh + n], &paras->ic[i * paras->ih + j]);

		l = 0;
		if(paras->rk == 0){
			l = 1;
		}
		if (m > l) {
			path = -1;
			MIN = INF;
			for (l = 0; l < 7; ++l) {
				x = IPX[l][0] + i;
				y = IPX[l][1] + j;
				if (x < 0 || y < 0 || y >= paras->ih) {
					continue;
				}
				ACCM *accm = get(paras->cl, m-1, x, y);
				int dx = accm->dxx + accm->dyy + accm->dxy;
				if (MIN > dx) {
					if (l == 6 && accm->x0 >= LIMIT) {
						continue;
					}
					MIN = dx;
					path = l;
				}
			}
			ACCM *accm = get(paras->cl, m, i, j);
			if (path < 0) {
				accm->dxx = ERR;
				accm->dxy = ERR;
			} else {
				x = IPX[path][0] + i;
				y = IPX[path][1] + j;
				ACCM *accm1 = get(paras->cl, m-1, x, y);
				accm->dxx = distance + accm1->dxx;
				accm->dxy = accm1->dxy + accm1->dyy;
				if (path == 6) {
					accm->x0 = accm1->x0 + 1;
				}
			}
		} else {	
			get(paras->cl, m, i, j)->dxx = distance;
		}
		if (n > 0) {
			path = -1;
			MIN = INF;
			for (l = 0; l < 7; ++l) {
				x = IPY[l][0] + i;
				y = IPY[l][1] + j;
				if (y < 0 || x < 0 || x >= paras->iw) {
					continue;
				}
				ACCM *accm = get(paras->pl, m, x, y);
				int dy = accm->dxx + accm->dyy + accm->dyx;
				if (MIN > dy) {
					if (l == 6 && accm->y0 >= LIMIT) {
						continue;
					}
					MIN = dy;
					path = l;
				}
			}
			ACCM *accm = get(paras->cl, m, i, j);
			if (path < 0) {
				accm->dyx = ERR;
				accm->dyy = ERR;
			} else {
				x = IPY[path][0] + i;
				y = IPY[path][1] + j;
				ACCM *accm1 = get(paras->pl, m, x, y);
				accm->dyy = distance + accm1->dyy;
				accm->dyx = accm1->dyx + accm1->dxx;
				if (path == 6) {
					accm->y0 = accm1->y0 + 1;
				}
			}
		}else {
			get(paras->cl, m, i, j)->dyy = distance;
		}
		cur += TOTAL;
	}
}

static color* copyImageToDevice(Image *image){
	color* dist;
	int k = sizeof(color) * image->width * image->height;
	CUDA_CHECK(cudaMalloc((void**)&dist, k));
	for(k = 0; k < image->width; ++k){
		CUDA_CHECK(cudaMemcpy(&dist[k * image->height], image->data[k],
				image->height * sizeof(color), cudaMemcpyHostToDevice));
	}
	return dist;
}

static __device__ short calDistance(color *c1, color *c2) {
	if(c1->b == 255 && c1->r == 0 && c1->g == 0){
		return DMAX;
	}else{
		return (abs(c1->r - c2->r) + abs(c1->g - c2->g) + abs(c1->b - c2->b));
	}
}

static __device__ ACCM * get(ACCM *cur, int m, int i, int j){
	return &cur[(m * paras->iw + i) * paras->ih + j];
}
